//
/*   core functions for Matrix * Vector operation program
*    Copyright (C) 2009 Goffredo Marocchi
*
*    This program is free software; you can redistribute it and/or modify
*    it under the terms of the GNU General Public License as published by
*    the Free Software Foundation; either version 2 of the License, or
*    (at your option) any later version.
*
*    This program is distributed in the hope that it will be useful,
*    but WITHOUT ANY WARRANTY; without even the implied warranty of
*    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
*    GNU General Public License for more details.
*
*    You should have received a copy of the GNU General Public License along
*    with this program; if not, write to the Free Software Foundation, Inc.,
*    51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
*/

/*
* Main program
*
* This program aims to test a standard Matrix * Vector multiplication. The matrix is square and NxN while the vector is Nx1.
* The Matrix is considered as column major and stored as a 1-D array properly indexed through a custom macro
* defined inside the IGGS_MatTest_defines.h file.
*/


#include "MatTest_helper.h"
#include "IGGS_CUDA_helper.h"
#include "cublas.h"

texture <float, 1, cudaReadModeElementType> texVecB; /**< Texture reference object*/

#include "MatTest_kernel.cu"

/** \brief Set-up and launch  the M*V CUDA kernel after invocation by the Fortran program.\n
 * The purpose of this function is to set-up and launch the custom written M*V CUDA kernel.
*
*	\param matA reference to the source matrix (RAM/System memory),
*			float*.
*	\param vecB reference to the source vector (RAM/System memory),
*			float*.
*	\param h_C_GPU reference to the result vector (RAM/System memory),
*			float*.
*	\return void
*	\sa STEPS, ROWS, COLS, start_timer(), stop_timer(), mat_vec()
*/
extern "C" void cudafunction_cuda_ (float* matA, float* vecB, float* h_C_GPU) {

    //Let's declare and initialize the timers for Host-Device transfers and kernel execution phase
    unsigned int timer1 = 0; float timer1_ms = 0.0f; //Total execution time timer (ms)
	start_timer(&timer1);//Total execution time Timer started

	float* d_A,* d_B,* d_C;
	d_A = null; d_B = null; 	d_C = null;

	AllocateHostMemory(&d_A, 20);
	FreeHostMemory(d_A);

#ifdef _DEBUG
	unsigned int timer_toGDDR = 0; float t_toGDDR_ms = 0.0f; //(main RAM --> PCI-Express --> GPU RAM)
	unsigned int timer_toRAM = 0; float t_toRAM_ms = 0.0f; //(GPU RAM --> PCI-Express --> main RAM)
	gpu_print_free_memory_(); send_msg("...allocating GPU memory.");
#endif

	Allocate_GPU_memory(d_A, DATA_M);
	Allocate_GPU_memory(d_B, DATA_V);
	Allocate_GPU_memory(d_C, DATA_V);

	SafeThreadSync("Memory allocated"); //let's make sure we wait until the data transfer is fully complete

#ifdef _DEBUG
	start_timer(&timer_toGDDR); send_msg("...copying input data to GPU mem.");
#endif

	Write_to_GPU(d_A, matA, DATA_M);
	Write_to_GPU(d_B, vecB, DATA_V);

	BindTex1D_rmetFP(NULL, texVecB, d_B); //we bind the source vector to a texture so we can read
																   //its values from the texture cache.
#ifdef _DEBUG
	stop_timer(timer_toGDDR, &t_toGDDR_ms); send_msg("Data init done.");
	#if SHARED_MEM == 1
		send_msg ("\n\"Using Shared Memory...\"\n");
	#endif
#endif

	SafeThreadSync("(check) CUDA driver API");

	dim3 dimBlock (TBLOCK);
	dim3 dimGrid (ROWS/TBLOCK); //1 Thread Block per ogni BLOCK_SIZE^2 righe di A
													 //ROWS is a multiple of TBLOCK

	MatTest<<<dimGrid, dimBlock>>>(d_C, d_A, d_B); //project's custom CUDA kernel invocation

	//SafeCall(cudaMemset((void *)d_C, 0, DATA_V));
	SafeThreadSync("(check) Kernel launch");

#ifdef _DEBUG
	start_timer(&timer_toRAM); send_msg("Reading back GPU result...\n");
#endif

	Read_from_GPU(h_C_GPU, d_C, DATA_V);

#ifdef _DEBUG
	stop_timer(timer_toRAM, &t_toRAM_ms);
#endif

	//Clean-up, freeing, of memory resources... the texture resource and the three arrays stored on the device...
    UnbindTex1D_rmetFP(texVecB);

    Free_GPU_memory(d_A);
    Free_GPU_memory(d_B);
    Free_GPU_memory(d_C);

    //PRINT_STEP();
	stop_timer(timer1, &timer1_ms);//Timer stopped
	send_msg ("Transfer + Exec + Readback time on GPU with CUDA: %f (ms)", timer1_ms);

#ifdef _DEBUG
	float exec_time_CUDA = (timer1_ms - t_toGDDR_ms - t_toRAM_ms);
	send_msg ("Execution time on GPU with CUDA: %f (ms)", exec_time_CUDA);
	send_msg ("Transfer to GPU with CUDA: %f (ms)", t_toGDDR_ms);
	send_msg ("Transfer from GPU with CUDA: %f (ms)\n", t_toRAM_ms);
	send_msg ("Risultati GPU (CUDA):");
	send_msg ("C_GPU.x= %f C_GPU.y= %f C_GPU.z= %f C_GPU.w= %f\n", h_C_GPU[0], h_C_GPU[1],
		h_C_GPU[2], h_C_GPU[3]);
	send_msg("Shutting down...\n");
#endif

	return;

}

/** \brief Function prepared to be called by the main Fortran function.\n
 * Its purpose is to set-up and execute the M*V operation using CUBLAS functions.
*
*	\param matA reference to the source matrix (RAM/System memory),
*			float*.
*	\param vecB reference to the source vector (RAM/System memory),
*			float*.
*	\param h_C_CUBLAS reference to the result vector (RAM/System memory),
*			float*.
*	\return void
*	\sa STEPS, ROWS, COLS, start_timer(), stop_timer(), mat_vec()
*/

extern "C" void cudafunction_cublas_ (float* matA, float* vecB, float* h_C_CUBLAS) {

	unsigned int timer2 = 0;
	float timer2_ms = 0.0f;
	start_timer(&timer2);

	float *d_A, *d_B, *d_C;

	d_A = null;
	d_B = null;
	d_C = null;

	CUBLAS_Init();

#ifdef _DEBUG
	unsigned int timer_toGDDR = 0; float t_toGDDR_ms = 0.0f;
	unsigned int timer_toRAM = 0; float t_toRAM_ms = 0.0f;
	send_msg("Allocating memory...\n");
#endif

	CUBLAS_Allocate_GPU_memory(d_A, ROWS*COLS);
	CUBLAS_Allocate_GPU_memory(d_B, ROWS);
	CUBLAS_Allocate_GPU_memory(d_C, ROWS);

	SafeThreadSync("d_A, d_B, and d_C have been allocated and will be transferred...");

#ifdef _DEBUG
	start_timer(&timer_toGDDR);
#endif

	CUBLAS_Write_to_GPU(d_A, matA, ROWS*COLS);
	CUBLAS_Write_to_GPU(d_B, vecB, ROWS);

#ifdef _DEBUG
	stop_timer(timer_toGDDR, &t_toGDDR_ms);
#endif

	SafeThreadSync("launching cublasSgemv()");
	cublasSgemv('n', ROWS, COLS, 1, d_A, ROWS, d_B, 1, 0, d_C, 1);
	cudaThreadSynchronize();
	CUBLAS_check_error ();

#ifdef _DEBUG
	start_timer(&timer_toRAM); send_msg("Reading back CUBLAS result...\n");
#endif

	// GPU --> main RAM (result readback)
	CUBLAS_Read_from_GPU(h_C_CUBLAS, d_C, ROWS);

#ifdef _DEBUG
	stop_timer(timer_toRAM, &t_toRAM_ms);
#endif

    CUBLAS_Free_GPU_memory(d_A);
    CUBLAS_Free_GPU_memory(d_B);
    CUBLAS_Free_GPU_memory(d_C);

	stop_timer(timer2, &timer2_ms);//Timer stopped
	send_msg ("Transfer + Exec + Readback time on GPU with CUBLAS: %f (ms)", timer2_ms);

#ifdef _DEBUG
	float exec_time_CUBLAS =  (timer2_ms - t_toGDDR_ms - t_toRAM_ms);
	send_msg ("Execution time on GPU with CUBLAS: %f (ms)", exec_time_CUBLAS);
	send_msg ("Transfer to GPU with CUBLAS: %f (ms)", t_toGDDR_ms);
	send_msg ("Transfer from GPU with CUBLAS: %f (ms)", t_toRAM_ms);
	send_msg ("\nRisultati GPU (CUBLAS):");
	send_msg ("C_GPU.x=%f C_GPU.y=%f C_GPU.z=%f C_GPU.w=%f\n", h_C_CUBLAS[0], h_C_CUBLAS[1],
		h_C_CUBLAS[2], h_C_CUBLAS[3]);
	send_msg("Shutting down...\n");
#endif

	return;

}


//
/*
* Copyright 1993-2007 NVIDIA Corporation.  All rights reserved.
*
* NOTICE TO USER:
*
* This source code is subject to NVIDIA ownership rights under U.S. and
* international Copyright laws.  Users and possessors of this source code
* are hereby granted a nonexclusive, royalty-free license to use this code
* in individual and commercial software.
*
* NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE
* CODE FOR ANY PURPOSE.  IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR
* IMPLIED WARRANTY OF ANY KIND.  NVIDIA DISCLAIMS ALL WARRANTIES WITH
* REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF
* MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
* IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL,
* OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS
* OF USE, DATA OR PROFITS,  WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE
* OR OTHER TORTIOUS ACTION,  ARISING OUT OF OR IN CONNECTION WITH THE USE
* OR PERFORMANCE OF THIS SOURCE CODE.
*
* U.S. Government End Users.   This source code is a "commercial item" as
* that term is defined at  48 C.F.R. 2.101 (OCT 1995), consisting  of
* "commercial computer  software"  and "commercial computer software
* documentation" as such terms are  used in 48 C.F.R. 12.212 (SEPT 1995)
* and is provided to the U.S. Government only as a commercial end item.
* Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through
* 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the
* source code with only those rights set forth herein.
*
* Any use of this source code in individual and commercial software must
* include, in the user documentation and internal comments to the code,
* the above Disclaimer and U.S. Government End Users Notice.
*/
